#include "kernel_caller.cuh"

void multiplyUsingKernels() {
	
	Kernel kernels[] = {
		func (multiplyKernel_8_noPrefetch_granularity_1x1_noLoopUnrolling, 8, 1),
		func (multiplyKernel_8_noPrefetch_granularity_1x1_2LoopUnrolling, 8, 1),
		func (multiplyKernel_8_noPrefetch_granularity_1x1_4LoopUnrolling, 8, 1),
		func (multiplyKernel_8_noPrefetch_granularity_1x1_fullLoopUnrolling, 8, 1),
		func (multiplyKernel_8_noPrefetch_granularity_1x2_noLoopUnrolling, 8, 2),
		func (multiplyKernel_8_noPrefetch_granularity_1x2_2LoopUnrolling, 8, 2),
		func (multiplyKernel_8_noPrefetch_granularity_1x2_4LoopUnrolling, 8, 2),
		func (multiplyKernel_8_noPrefetch_granularity_1x2_fullLoopUnrolling, 8, 2),
		func (multiplyKernel_8_noPrefetch_granularity_1x4_noLoopUnrolling, 8, 4),
		func (multiplyKernel_8_noPrefetch_granularity_1x4_2LoopUnrolling, 8, 4),
		func (multiplyKernel_8_noPrefetch_granularity_1x4_4LoopUnrolling, 8, 4),
		func (multiplyKernel_8_noPrefetch_granularity_1x4_fullLoopUnrolling, 8, 4),
		func (multiplyKernel_8_prefetch_granularity_1x1_noLoopUnrolling, 8, 1),
		func (multiplyKernel_8_prefetch_granularity_1x1_2LoopUnrolling, 8, 1),
		func (multiplyKernel_8_prefetch_granularity_1x1_4LoopUnrolling, 8, 1),
		func (multiplyKernel_8_prefetch_granularity_1x1_fullLoopUnrolling, 8, 1),
		func (multiplyKernel_8_prefetch_granularity_1x2_noLoopUnrolling, 8, 2),
		func (multiplyKernel_8_prefetch_granularity_1x2_2LoopUnrolling, 8, 2),
		func (multiplyKernel_8_prefetch_granularity_1x2_4LoopUnrolling, 8, 2),
		func (multiplyKernel_8_prefetch_granularity_1x2_fullLoopUnrolling, 8, 2),
		func (multiplyKernel_8_prefetch_granularity_1x4_noLoopUnrolling, 8, 4),
		func (multiplyKernel_8_prefetch_granularity_1x4_2LoopUnrolling, 8, 4),
		func (multiplyKernel_8_prefetch_granularity_1x4_4LoopUnrolling, 8, 4),
		func (multiplyKernel_8_prefetch_granularity_1x4_fullLoopUnrolling, 8, 4),
		func (multiplyKernel_16_noPrefetch_granularity_1x1_noLoopUnrolling, 16, 1),
		func (multiplyKernel_16_noPrefetch_granularity_1x1_2LoopUnrolling, 16, 1),
		func (multiplyKernel_16_noPrefetch_granularity_1x1_4LoopUnrolling, 16, 1),
		func (multiplyKernel_16_noPrefetch_granularity_1x1_fullLoopUnrolling, 16, 1),
		func (multiplyKernel_16_noPrefetch_granularity_1x2_noLoopUnrolling, 16, 2),
		func (multiplyKernel_16_noPrefetch_granularity_1x2_2LoopUnrolling, 16, 2),
		func (multiplyKernel_16_noPrefetch_granularity_1x2_4LoopUnrolling, 16, 2),
		func (multiplyKernel_16_noPrefetch_granularity_1x2_fullLoopUnrolling, 16, 2),
		func (multiplyKernel_16_noPrefetch_granularity_1x4_noLoopUnrolling, 16, 4),
		func (multiplyKernel_16_noPrefetch_granularity_1x4_2LoopUnrolling, 16, 4),
		func (multiplyKernel_16_noPrefetch_granularity_1x4_4LoopUnrolling, 16, 4),
		func (multiplyKernel_16_noPrefetch_granularity_1x4_fullLoopUnrolling, 16, 4),
		func (multiplyKernel_16_prefetch_granularity_1x1_noLoopUnrolling, 16, 1),
		func (multiplyKernel_16_prefetch_granularity_1x1_2LoopUnrolling, 16, 1),
		func (multiplyKernel_16_prefetch_granularity_1x1_4LoopUnrolling, 16, 1),
		func (multiplyKernel_16_prefetch_granularity_1x1_fullLoopUnrolling, 16, 1),
		func (multiplyKernel_16_prefetch_granularity_1x2_noLoopUnrolling, 16, 2),
		func (multiplyKernel_16_prefetch_granularity_1x2_2LoopUnrolling, 16, 2),
		func (multiplyKernel_16_prefetch_granularity_1x2_4LoopUnrolling, 16, 2),
		func (multiplyKernel_16_prefetch_granularity_1x2_fullLoopUnrolling, 16, 2),
		func (multiplyKernel_16_prefetch_granularity_1x4_noLoopUnrolling, 16, 4),
		func (multiplyKernel_16_prefetch_granularity_1x4_2LoopUnrolling, 16, 4),
		func (multiplyKernel_16_prefetch_granularity_1x4_4LoopUnrolling, 16, 4),
		func (multiplyKernel_16_prefetch_granularity_1x4_fullLoopUnrolling, 16, 4)
	};

	callKernels(kernels, 48);
}

void callKernels(Kernel *kernels, int kernelsCount) {
	cout << "Reading input.. ";
	float *matrix1, *matrix2, *correctResult;
	int height, length, width;
	readMatrixFromFile(matrix1, height, length, "matrix1.txt");
	readMatrixFromFile(matrix2, length, width, "matrix2.txt");
	readMatrixFromFile(correctResult, height, width, "result.txt");
	cout << "done." << endl;

	float *emptyMatrix = createZeroMatrix(height, width);

	float *dev_a = 0;
	float *dev_b = 0;
	float *dev_c = 0;
	cudaSetDevice(0);
    
	cudaMalloc((void**)&dev_a, height * length * sizeof(float));
	cudaMalloc((void**)&dev_b, length * width * sizeof(float));
    
	cudaMemcpy(dev_a, matrix1, height * length * sizeof(float), cudaMemcpyHostToDevice);
	cudaMemcpy(dev_b, matrix2, length * width * sizeof(float), cudaMemcpyHostToDevice);
 
	int successfulRunsCount = 0;
	float operationsCount = 2.0f * height * width * length;

	clock_t time;
	
	cout << "Running Kernels." << endl;
	cout << endl;
	
	for (int i = 0; i < kernelsCount; i++) {
		cout << "Test # " << (i + 1) << ":" << endl;

		Function function = kernels[i].first.first;
		string functionName = kernels[i].first.second;
		int tileWidth = kernels[i].second.first;
		int granularity = kernels[i].second.second;
		
		int tilesCountX = 1 + (width - 1) / (tileWidth * granularity);
		int tilesCountY = 1 + (height - 1) / tileWidth;

		dim3 dimGrid(tilesCountX, tilesCountY);
		dim3 dimBlock(tileWidth, tileWidth);

		cudaMalloc((void**)&dev_c, width * height * sizeof(float));
		cudaMemcpy(dev_c, emptyMatrix, width * height * sizeof(float), cudaMemcpyHostToDevice);

		cout << "Now Running " << functionName.c_str() << endl;

		function<<<dimGrid, dimBlock>>>(dev_c, dev_a, dev_b, height, width, length);

		time = clock();
		cudaDeviceSynchronize();
		time = clock() - time;
		cout << "Finished in " << time << " milliseconds." << endl;

		float GFLOPS = 1e-6 * operationsCount / time;
		cout << "GFLOPS = " << GFLOPS << "." << endl;

		float *result = new float[width * height];
		cudaMemcpy(result, dev_c, width * height * sizeof(float), cudaMemcpyDeviceToHost);

		cudaFree(dev_c);

		cout << "Checking result.. ";
		bool verification = verify(result, correctResult, height, width);
		successfulRunsCount += verification;

		if (verification)
			cout << "correct." << endl;
		else
			cout << "incorrect." << endl;
			
		cout << endl;
	}

	cout << "Runs completed, " << successfulRunsCount << "/" << kernelsCount << " succeeded." << endl;
	cout << "Success rate : " << (100.0f * successfulRunsCount / kernelsCount) << "%." << endl;
	cudaFree(dev_a);
	cudaFree(dev_b);

    cudaDeviceReset();
}